{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "name": "udacity-cs344-hw2",
      "version": "0.3.2",
      "provenance": [],
      "collapsed_sections": [],
      "include_colab_link": true
    },
    "kernelspec": {
      "name": "python3",
      "display_name": "Python 3"
    },
    "accelerator": "GPU"
  },
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "view-in-github",
        "colab_type": "text"
      },
      "source": [
        "<a href=\"https://colab.research.google.com/github/depctg/udacity-cs344-colab/blob/master/notebook/udacity_cs344_hw2.ipynb\" target=\"_parent\"><img src=\"https://colab.research.google.com/assets/colab-badge.svg\" alt=\"Open In Colab\"/></a>"
      ]
    },
    {
      "metadata": {
        "id": "hse6gSyUS5ka",
        "colab_type": "code",
        "colab": {}
      },
      "cell_type": "code",
      "source": [
        "# Homework 2 for Udacity CS344 Course, Intro to Parallel Programming\n",
        "# clone the code repo,\n",
        "!git clone https://github.com/depctg/udacity-cs344-colab\n",
        "!pip install git+git://github.com/depctg/nvcc4jupyter.git\n",
        "\n",
        "# load cuda plugin\n",
        "%config NVCCPluginV2.static_dir = True\n",
        "%config NVCCPluginV2.relative_dir = \"udacity-cs344-colab/src/HW2\"\n",
        "%load_ext nvcc_plugin\n",
        "\n",
        "# change to work directory, generate makefiles\n",
        "!mkdir udacity-cs344-colab/build\n",
        "%cd udacity-cs344-colab/build\n",
        "!cmake ../src"
      ],
      "execution_count": 0,
      "outputs": []
    },
    {
      "metadata": {
        "id": "3vA0JP15TORh",
        "colab_type": "code",
        "colab": {}
      },
      "cell_type": "code",
      "source": [
        "%%cuda --name student_func.cu\n",
        "\n",
        "// Homework 2\n",
        "// Image Blurring\n",
        "//\n",
        "// In this homework we are blurring an image. To do this, imagine that we have\n",
        "// a square array of weight values. For each pixel in the image, imagine that we\n",
        "// overlay this square array of weights on top of the image such that the center\n",
        "// of the weight array is aligned with the current pixel. To compute a blurred\n",
        "// pixel value, we multiply each pair of numbers that line up. In other words, we\n",
        "// multiply each weight with the pixel underneath it. Finally, we add up all of the\n",
        "// multiplied numbers and assign that value to our output for the current pixel.\n",
        "// We repeat this process for all the pixels in the image.\n",
        "\n",
        "// To help get you started, we have included some useful notes here.\n",
        "\n",
        "//****************************************************************************\n",
        "\n",
        "// For a color image that has multiple channels, we suggest separating\n",
        "// the different color channels so that each color is stored contiguously\n",
        "// instead of being interleaved. This will simplify your code.\n",
        "\n",
        "// That is instead of RGBARGBARGBARGBA... we suggest transforming to three\n",
        "// arrays (as in the previous homework we ignore the alpha channel again):\n",
        "//  1) RRRRRRRR...\n",
        "//  2) GGGGGGGG...\n",
        "//  3) BBBBBBBB...\n",
        "//\n",
        "// The original layout is known an Array of Structures (AoS) whereas the\n",
        "// format we are converting to is known as a Structure of Arrays (SoA).\n",
        "\n",
        "// As a warm-up, we will ask you to write the kernel that performs this\n",
        "// separation. You should then write the \"meat\" of the assignment,\n",
        "// which is the kernel that performs the actual blur. We provide code that\n",
        "// re-combines your blurred results for each color channel.\n",
        "\n",
        "//****************************************************************************\n",
        "\n",
        "// You must fill in the gaussian_blur kernel to perform the blurring of the\n",
        "// inputChannel, using the array of weights, and put the result in the outputChannel.\n",
        "\n",
        "// Here is an example of computing a blur, using a weighted average, for a single\n",
        "// pixel in a small image.\n",
        "//\n",
        "// Array of weights:\n",
        "//\n",
        "//  0.0  0.2  0.0\n",
        "//  0.2  0.2  0.2\n",
        "//  0.0  0.2  0.0\n",
        "//\n",
        "// Image (note that we align the array of weights to the center of the box):\n",
        "//\n",
        "//    1  2  5  2  0  3\n",
        "//       -------\n",
        "//    3 |2  5  1| 6  0       0.0*2 + 0.2*5 + 0.0*1 +\n",
        "//      |       |\n",
        "//    4 |3  6  2| 1  4   ->  0.2*3 + 0.2*6 + 0.2*2 +   ->  3.2\n",
        "//      |       |\n",
        "//    0 |4  0  3| 4  2       0.0*4 + 0.2*0 + 0.0*3\n",
        "//       -------\n",
        "//    9  6  5  0  3  9\n",
        "//\n",
        "//         (1)                         (2)                 (3)\n",
        "//\n",
        "// A good starting place is to map each thread to a pixel as you have before.\n",
        "// Then every thread can perform steps 2 and 3 in the diagram above\n",
        "// completely independently of one another.\n",
        "\n",
        "// Note that the array of weights is square, so its height is the same as its width.\n",
        "// We refer to the array of weights as a filter, and we refer to its width with the\n",
        "// variable filterWidth.\n",
        "\n",
        "//****************************************************************************\n",
        "\n",
        "// Your homework submission will be evaluated based on correctness and speed.\n",
        "// We test each pixel against a reference solution. If any pixel differs by\n",
        "// more than some small threshold value, the system will tell you that your\n",
        "// solution is incorrect, and it will let you try again.\n",
        "\n",
        "// Once you have gotten that working correctly, then you can think about using\n",
        "// shared memory and having the threads cooperate to achieve better performance.\n",
        "\n",
        "//****************************************************************************\n",
        "\n",
        "// Also note that we've supplied a helpful debugging function called checkCudaErrors.\n",
        "// You should wrap your allocation and copying statements like we've done in the\n",
        "// code we're supplying you. Here is an example of the unsafe way to allocate\n",
        "// memory on the GPU:\n",
        "//\n",
        "// cudaMalloc(&d_red, sizeof(unsigned char) * numRows * numCols);\n",
        "//\n",
        "// Here is an example of the safe way to do the same thing:\n",
        "//\n",
        "// checkCudaErrors(cudaMalloc(&d_red, sizeof(unsigned char) * numRows * numCols));\n",
        "//\n",
        "// Writing code the safe way requires slightly more typing, but is very helpful for\n",
        "// catching mistakes. If you write code the unsafe way and you make a mistake, then\n",
        "// any subsequent kernels won't compute anything, and it will be hard to figure out\n",
        "// why. Writing code the safe way will inform you as soon as you make a mistake.\n",
        "\n",
        "// Finally, remember to free the memory you allocate at the end of the function.\n",
        "\n",
        "//****************************************************************************\n",
        "\n",
        "#include \"utils.h\"\n",
        "\n",
        "__global__\n",
        "void gaussian_blur(const unsigned char* const inputChannel,\n",
        "                   unsigned char* const outputChannel,\n",
        "                   int numRows, int numCols,\n",
        "                   const float* const filter, const int filterWidth)\n",
        "{\n",
        "  // TODO\n",
        "\n",
        "  // NOTE: Be sure to compute any intermediate results in floating point\n",
        "  // before storing the final result as unsigned char.\n",
        "\n",
        "  // NOTE: Be careful not to try to access memory that is outside the bounds of\n",
        "  // the image. You'll want code that performs the following check before accessing\n",
        "  // GPU memory:\n",
        "  //\n",
        "  // if ( absolute_image_position_x >= numCols ||\n",
        "  //      absolute_image_position_y >= numRows )\n",
        "  // {\n",
        "  //     return;\n",
        "  // }\n",
        "\n",
        "  // NOTE: If a thread's absolute position 2D position is within the image, but some of\n",
        "  // its neighbors are outside the image, then you will need to be extra careful. Instead\n",
        "  // of trying to read such a neighbor value from GPU memory (which won't work because\n",
        "  // the value is out of bounds), you should explicitly clamp the neighbor values you read\n",
        "  // to be within the bounds of the image. If this is not clear to you, then please refer\n",
        "  // to sequential reference solution for the exact clamping semantics you should follow.\n",
        "}\n",
        "\n",
        "//This kernel takes in an image represented as a uchar4 and splits\n",
        "//it into three images consisting of only one color channel each\n",
        "__global__\n",
        "void separateChannels(const uchar4* const inputImageRGBA,\n",
        "                      int numRows,\n",
        "                      int numCols,\n",
        "                      unsigned char* const redChannel,\n",
        "                      unsigned char* const greenChannel,\n",
        "                      unsigned char* const blueChannel)\n",
        "{\n",
        "  // TODO\n",
        "  //\n",
        "  // NOTE: Be careful not to try to access memory that is outside the bounds of\n",
        "  // the image. You'll want code that performs the following check before accessing\n",
        "  // GPU memory:\n",
        "  //\n",
        "  // if ( absolute_image_position_x >= numCols ||\n",
        "  //      absolute_image_position_y >= numRows )\n",
        "  // {\n",
        "  //     return;\n",
        "  // }\n",
        "}\n",
        "\n",
        "//This kernel takes in three color channels and recombines them\n",
        "//into one image.  The alpha channel is set to 255 to represent\n",
        "//that this image has no transparency.\n",
        "__global__\n",
        "void recombineChannels(const unsigned char* const redChannel,\n",
        "                       const unsigned char* const greenChannel,\n",
        "                       const unsigned char* const blueChannel,\n",
        "                       uchar4* const outputImageRGBA,\n",
        "                       int numRows,\n",
        "                       int numCols)\n",
        "{\n",
        "  const int2 thread_2D_pos = make_int2( blockIdx.x * blockDim.x + threadIdx.x,\n",
        "                                        blockIdx.y * blockDim.y + threadIdx.y);\n",
        "\n",
        "  const int thread_1D_pos = thread_2D_pos.y * numCols + thread_2D_pos.x;\n",
        "\n",
        "  //make sure we don't try and access memory outside the image\n",
        "  //by having any threads mapped there return early\n",
        "  if (thread_2D_pos.x >= numCols || thread_2D_pos.y >= numRows)\n",
        "    return;\n",
        "\n",
        "  unsigned char red   = redChannel[thread_1D_pos];\n",
        "  unsigned char green = greenChannel[thread_1D_pos];\n",
        "  unsigned char blue  = blueChannel[thread_1D_pos];\n",
        "\n",
        "  //Alpha should be 255 for no transparency\n",
        "  uchar4 outputPixel = make_uchar4(red, green, blue, 255);\n",
        "\n",
        "  outputImageRGBA[thread_1D_pos] = outputPixel;\n",
        "}\n",
        "\n",
        "unsigned char *d_red, *d_green, *d_blue;\n",
        "float         *d_filter;\n",
        "\n",
        "void allocateMemoryAndCopyToGPU(const size_t numRowsImage, const size_t numColsImage,\n",
        "                                const float* const h_filter, const size_t filterWidth)\n",
        "{\n",
        "\n",
        "  //allocate memory for the three different channels\n",
        "  //original\n",
        "  checkCudaErrors(cudaMalloc(&d_red,   sizeof(unsigned char) * numRowsImage * numColsImage));\n",
        "  checkCudaErrors(cudaMalloc(&d_green, sizeof(unsigned char) * numRowsImage * numColsImage));\n",
        "  checkCudaErrors(cudaMalloc(&d_blue,  sizeof(unsigned char) * numRowsImage * numColsImage));\n",
        "\n",
        "  //TODO:\n",
        "  //Allocate memory for the filter on the GPU\n",
        "  //Use the pointer d_filter that we have already declared for you\n",
        "  //You need to allocate memory for the filter with cudaMalloc\n",
        "  //be sure to use checkCudaErrors like the above examples to\n",
        "  //be able to tell if anything goes wrong\n",
        "  //IMPORTANT: Notice that we pass a pointer to a pointer to cudaMalloc\n",
        "\n",
        "  //TODO:\n",
        "  //Copy the filter on the host (h_filter) to the memory you just allocated\n",
        "  //on the GPU.  cudaMemcpy(dst, src, numBytes, cudaMemcpyHostToDevice);\n",
        "  //Remember to use checkCudaErrors!\n",
        "\n",
        "}\n",
        "\n",
        "void your_gaussian_blur(const uchar4 * const h_inputImageRGBA, uchar4 * const d_inputImageRGBA,\n",
        "                        uchar4* const d_outputImageRGBA, const size_t numRows, const size_t numCols,\n",
        "                        unsigned char *d_redBlurred,\n",
        "                        unsigned char *d_greenBlurred,\n",
        "                        unsigned char *d_blueBlurred,\n",
        "                        const int filterWidth)\n",
        "{\n",
        "  //TODO: Set reasonable block size (i.e., number of threads per block)\n",
        "  const dim3 blockSize;\n",
        "\n",
        "  //TODO:\n",
        "  //Compute correct grid size (i.e., number of blocks per kernel launch)\n",
        "  //from the image size and and block size.\n",
        "  const dim3 gridSize;\n",
        "\n",
        "  //TODO: Launch a kernel for separating the RGBA image into different color channels\n",
        "\n",
        "  // Call cudaDeviceSynchronize(), then call checkCudaErrors() immediately after\n",
        "  // launching your kernel to make sure that you didn't make any mistakes.\n",
        "  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());\n",
        "\n",
        "  //TODO: Call your convolution kernel here 3 times, once for each color channel.\n",
        "\n",
        "  // Again, call cudaDeviceSynchronize(), then call checkCudaErrors() immediately after\n",
        "  // launching your kernel to make sure that you didn't make any mistakes.\n",
        "  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());\n",
        "\n",
        "  // Now we recombine your results. We take care of launching this kernel for you.\n",
        "  //\n",
        "  // NOTE: This kernel launch depends on the gridSize and blockSize variables,\n",
        "  // which you must set yourself.\n",
        "  recombineChannels<<<gridSize, blockSize>>>(d_redBlurred,\n",
        "                                             d_greenBlurred,\n",
        "                                             d_blueBlurred,\n",
        "                                             d_outputImageRGBA,\n",
        "                                             numRows,\n",
        "                                             numCols);\n",
        "  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());\n",
        "\n",
        "}\n",
        "\n",
        "\n",
        "//Free all the memory that we allocated\n",
        "//TODO: make sure you free any arrays that you allocated\n",
        "void cleanup() {\n",
        "  checkCudaErrors(cudaFree(d_red));\n",
        "  checkCudaErrors(cudaFree(d_green));\n",
        "  checkCudaErrors(cudaFree(d_blue));\n",
        "}"
      ],
      "execution_count": 0,
      "outputs": []
    },
    {
      "metadata": {
        "id": "sSAnpiE2nL1T",
        "colab_type": "code",
        "colab": {}
      },
      "cell_type": "code",
      "source": [
        "# make the cuda project\n",
        "!make HW2\n",
        "print(\"\\n====== RESULT OF HW2 =======\\n\")\n",
        "!bin/HW2 ../src/HW1/cinque_terre.gold"
      ],
      "execution_count": 0,
      "outputs": []
    },
    {
      "metadata": {
        "id": "2cGKiWj_n-Na",
        "colab_type": "code",
        "colab": {}
      },
      "cell_type": "code",
      "source": [
        "# plot output images\n",
        "import matplotlib.pyplot as plt\n",
        "_,ax = plt.subplots(2,2, dpi=150)\n",
        "\n",
        "ax[0][0].imshow(plt.imread(\"../src/HW1/cinque_terre_small.jpg\"))\n",
        "ax[0][0].set_title(\"original\")\n",
        "ax[0][0].grid(False)\n",
        "\n",
        "ax[0][1].imshow(plt.imread(\"HW2_output.png\"))\n",
        "ax[0][1].set_title(\"output\")\n",
        "ax[0][1].grid(False)\n",
        "\n",
        "ax[1][0].imshow(plt.imread(\"HW2_reference.png\"))\n",
        "ax[1][0].set_title(\"reference\")\n",
        "ax[1][0].grid(False)\n",
        "\n",
        "ax[1][1].imshow(plt.imread(\"HW2_differenceImage.png\"))\n",
        "ax[1][1].set_title(\"difference\")\n",
        "ax[1][1].grid(False)\n",
        "\n",
        "plt.show()"
      ],
      "execution_count": 0,
      "outputs": []
    }
  ]
}